// GRUPO 02 - TURMA A
// Gabriel Tessaroli   – 4321350
// Lucas Nunes Arvani  – 7986802
// Lucas Fúrio Franco  – 6793010

#include <cuda.h>
#include <sys/time.h>
#include <stdio.h>
#include <stdlib.h>

__global__ void smooth(unsigned char *image, unsigned char *image_smooth, int N, int n_total, int cols, int rows);
void write_ppm(unsigned char* image_smooth, int cols, int rows, int max);
int p(int r, int c, int p, int cols);
int open_ppm(unsigned char **img ,int *cols, int *rows, int *max);
void print_time(struct timeval begin, struct timeval end, int flag) ;

int main(int argc, char *argv[]){
  int param = 5;

  struct timeval b_prog, e_prog;
  struct timeval b_proc, e_proc;

  gettimeofday(&b_prog, NULL);

  int cols, rows, max, vector_size;
  // Vetor de dados da imagem a ser processada
  unsigned char *img, *img_d;
  // Vetor com valores do resultado da imagem
  unsigned char *img_smooth, *img_smooth_d;

  // atribui valores para 'cols', 'rows' e 'max'
  // e popula vetor de dados da imagem
  open_ppm(&img, &cols, &rows, &max);
  vector_size = cols*rows*3*sizeof(unsigned char);

  // Reserva memória para imagem resultante
  img_smooth = (unsigned char*) malloc(vector_size);

  gettimeofday(&b_proc, NULL);

  // Reserva memória da placa de video para receber as imagens
  cudaMalloc(&img_d, vector_size);
  cudaMalloc(&img_smooth_d, vector_size);

  // transfere a imagem para a placa, onde será executado o algoritmo
  cudaMemcpy(img_d, img, vector_size, cudaMemcpyHostToDevice);

  // algoritmo é realizado usando (tam_vetor/256+1) blocos e 256 threads/bloco
  smooth<<<ceil(vector_size/ 512 + 1), 512>>>(img_d, img_smooth_d, param, vector_size, cols, rows);

  // transfere resultado de volta para host
  cudaMemcpy(img_smooth, img_smooth_d, vector_size, cudaMemcpyDeviceToHost);

  gettimeofday(&e_proc, NULL);

  // escreve imagem em um arquivo ppm
  write_ppm(img_smooth, cols, rows, max);

  gettimeofday(&e_prog, NULL);

  print_time(b_prog, e_prog, 1);
  print_time(b_proc, e_proc, 0);

  // Libera memória do device e host
  cudaFree(img_d);
  cudaFree(img_smooth_d);
  free(img);
  free(img_smooth);

  return 0;
}

// Função q processa a imagem
__global__ void smooth(unsigned char *image, unsigned char *image_smooth, int N, int n_total, int cols, int rows) {

  int sum;
  int avg;
  int border = N / 2;
  int c, r;
  int channel;

  // Atribui um 'id' que será unico para a thread
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  //obtem linha e coluna designada aquela thread
  r = i / cols;
  c = i % cols;

  // verefica se a thread é valida e realiza a paralelização
  if(i < (cols*rows)){

    // realiza o calculo do smooth no pixel
    for(channel = 0; channel < 3; ++channel) {
      sum = 0;
	  //N x N
      for(int a = (r-border); a <= (r+border); ++a)
        for(int b = (c-border); b <= (c+border); ++b)
          sum += (a>=0&&a<rows&&b>=0&&b<cols)? image[(((a*cols)+b)*3)+channel] : 0;

      avg = sum / (N*N);
      image_smooth[((((r*cols)+c)*3)+channel)] = avg;
	}
  }

}

// escreve imagem final em um arquivo
void write_ppm(unsigned char* image_smooth, int cols, int rows, int max) {
  FILE *out = fopen("out.ppm", "w");

  fprintf(out, "P3\n");
  fprintf(out, "%d %d %d", cols, rows, max);
  int i;

  for(i = 0; i < cols*rows*3; ++i) {
	if (!(i % 18)) fprintf(out, "\n");
    fprintf(out, "%d ", image_smooth[i]);
  }

  fclose(out);
}

// abre imagem e guarda seus dados em um vetor.
int open_ppm(unsigned char **img ,int *cols, int *rows, int *max) {
  char key[4];
  scanf("%s", key);
  if (strcmp(key,"P3")) {
    return -1;
  }

  scanf("%d %d %d\n", cols, rows, max);

  (*img) = (unsigned char *) malloc(sizeof(unsigned char)*(*cols)*(*rows)*3);

  int i;
  int buffer;
  for(i = 0; i < (*cols)*(*rows)*3; ++i) {
    scanf("%d ", &buffer);
	(*img)[i] = buffer;
  }

  return 1;
}

// calculo do tempo
void print_time(struct timeval begin, struct timeval end, int flag) {
	if (flag) {
		printf("Program Time: ");
	} else {
		printf("Execution time: ");
	}

	if (end.tv_usec < begin.tv_usec) {
		printf("%ld.%06lds.\n", end.tv_sec - begin.tv_sec - 1, (-1)*(end.tv_usec - begin.tv_usec));
	} else {
		printf("%ld.%06lds.\n", end.tv_sec - begin.tv_sec, (end.tv_usec - begin.tv_usec));
	}
}
